//#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
//#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
//#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable
//#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable

//#pragma OPENCL EXTENSION cl_khr_int64_base_atomics: enable
//#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics: enable ///nvidia doesnt support this
//#pragma OPENCL EXTENSION CL_KHR_gl_sharing : enable

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable

int min3(float x, float y, float z)
{

    return min(min(x,y),z);

}

int max3( float x,  float y,  float z)
{

    return max(max(x,y),z);

}

float rabs(float a)
{
    if(a<0)
    {
        return -a;
    }
    return a;
}

int out_of_bounds(float val, float min, float max)
{
    if(val > min && val < max)
    {
        return false;
    }
    return true;

}

float fastdist2(float x, float y)
{
// fast pythagoras approximation
    if(x<0) x=-x;
    if(y<0) y=-y;

    if(x>y)
        return (x+(y/4.0f));
    else
        return (y+(x/4.0f));
//	return sqrt(x*x+y*y);
}

float fastdist3(float x, float y, float z)
{
// the same, only for 3 dimensions
    float res;

    if(x<0) x=-x;
    if(y<0) y=-y;
    if(z<0) z=-z;

    if(x>y)
        res = (x+(y/4.0f));
    else
        res = (y+(x/4.0f));
    return fastdist2(res, z);
//return sqrt(x*x+y*y+z*z);
}


float form(float x1, float y1, float x2, float y2, float x3, float y3, float x, float y, int which)
{

    if(which==1)
    {
        return fabs(((x2*y-x*y2)+(x3*y2-x2*y3)+(x*y3-x3*y))/2.0);
    }
    else if(which==2)
    {
        return fabs(((x*y1-x1*y)+(x3*y-x*y3)+(x1*y3-x3*y1))/2.0);
    }
    else if(which==3)
    {
        return fabs(((x2*y1-x1*y2)+(x*y2-x2*y)+(x1*y-x*y1))/2.0);
    }

    return fabs((x2*y1-x1*y2)+(x3*y2-x2*y3)+(x1*y3-x3*y1))/2.0;

} ///triangle equations


float dcalc(float value)
{
    float depth_far=7000;
    bool t=false;
    // if(value<0)
    //{
    //t=true;
    //   value=-value;

    //}

    value=(log(value + 1)/(log(depth_far + 1)));

    //if(t){
    //    value=-value;
    //}

    return value;
}

float idcalc(float value)
{
    float depth_far=7000;

    value=exp(value*log(depth_far + 1))-1;


    return value;

}



__constant float depth_cutoff=0.22f;
__constant uint mulint=UINT_MAX;

__kernel void part1(__global float4 *c_pos, __global float4 *c_rot,  __global unsigned int* array_width, __global unsigned int* array_height, __global float4 *stuff, __global uint *screena, __global float4 *lights, __global uint *lightnum, __global uint *ids, __global uint *tcount, __global uint *dcount, __global uint *idbuffer, __global bool * state)
{

    ///camera position, camera rotation, screen width, screen height, position/normals/colour(unused), screen r g b depth components
    unsigned int i = get_global_id(0);
    unsigned int size=get_global_size(0);

    int screenwidth=array_width[0];
    int screenheight=array_height[0];

    const sampler_t RAWRAWRAWRAWR=      CLK_NORMALIZED_COORDS_FALSE |
                                        CLK_ADDRESS_CLAMP_TO_EDGE   |
                                        CLK_FILTER_LINEAR;

    bool isexit=false;

    if(i<0)
    {
        return;
    }


    int cstat=4;
    if(!(*state)){ ///state is true when i want multiple triangles per thread
        cstat=1;
    }
    i*=cstat;

    for(int aid=0; aid<cstat; aid++, i++){

    isexit=false;

    if(i >= *tcount)
    {
        return;
    }

    float depth_far=70000000;

    unsigned int location=i*12;


    float4 position[3];
    position[0]=stuff[location + 0];
    position[1]=stuff[location + 1];
    position[2]=stuff[location + 2];

    float4 normal[3];
    normal[0]=stuff[location + 3];
    normal[1]=stuff[location + 4];
    normal[2]=stuff[location + 5];

    float4 colour[3];
    colour[0]=stuff[location + 6];
    colour[1]=stuff[location + 7];
    colour[2]=stuff[location + 8];

    float4 uvw[3];
    uvw[0]=stuff[location + 9];
    uvw[1]=stuff[location + 10];
    uvw[2]=stuff[location + 11];


    ///if(gidbuffer!=dontdraw)
    ids[i]=0;





    float4 one, two, three; ///should have used an array


    float prez1, prez2, prez3;



    //float depth_cutoff=0.1;

    //float opart1=

    float opart1=(native_cos(c_rot[0].y)*(position[0].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[0].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[0].x-c_pos[0].x)));
    float opart2=(native_cos(c_rot[0].z)*(position[0].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[0].x-c_pos[0].x));


    one.x=      native_cos(c_rot[0].y)*(native_sin(c_rot[0].z)+native_cos(c_rot[0].z)*(position[0].x-c_pos[0].x)) - native_sin(c_rot[0].y)*(position[0].z-c_pos[0].z);
    one.y=      native_sin(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[0].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[0].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[0].x-c_pos[0].x)))+native_cos(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[0].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[0].x-c_pos[0].x));
    prez1=one.z=native_cos(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[0].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[0].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[0].x-c_pos[0].x)))-native_sin(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[0].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[0].x-c_pos[0].x));

    float rx;
    rx=one.x * (700.0f/one.z);
    float ry;
    ry=one.y * (700.0f/one.z);

    rx+=array_width[0]/2.0f;
    ry+=array_height[0]/2.0f;


    one.x=rx;
    one.y=ry;
    one.z=dcalc(one.z);






    two.x=native_cos(c_rot[0].y)*(native_sin(c_rot[0].z)+native_cos(c_rot[0].z)*(position[1].x-c_pos[0].x)) - native_sin(c_rot[0].y)*(position[1].z-c_pos[0].z);
    two.y=      native_sin(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[1].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[1].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[1].x-c_pos[0].x)))+native_cos(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[1].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[1].x-c_pos[0].x));
    prez2=two.z=native_cos(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[1].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[1].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[1].x-c_pos[0].x)))-native_sin(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[1].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[1].x-c_pos[0].x));


    rx=two.x * (700.0f/two.z);
    ry=two.y * (700.0f/two.z);

    rx+=array_width[0]/2.0f;
    ry+=array_height[0]/2.0f;


    two.x=rx;
    two.y=ry;
    two.z=dcalc(two.z);





    three.x=native_cos(c_rot[0].y)*(native_sin(c_rot[0].z)+native_cos(c_rot[0].z)*(position[2].x-c_pos[0].x)) - native_sin(c_rot[0].y)*(position[2].z-c_pos[0].z);
    three.y=      native_sin(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[2].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[2].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[2].x-c_pos[0].x)))+native_cos(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[2].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[2].x-c_pos[0].x));
    prez3=three.z=native_cos(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[2].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[2].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[2].x-c_pos[0].x)))-native_sin(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[2].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[2].x-c_pos[0].x));


    rx=three.x * (700.0f/three.z);
    ry=three.y * (700.0f/three.z);

    rx+=array_width[0]/2.0f;
    ry+=array_height[0]/2.0f;


    three.x=rx;
    three.y=ry;
    three.z=dcalc(three.z);










    if((three.z < depth_cutoff && two.z < depth_cutoff && one.z < depth_cutoff) || isnan(three.z) || isnan(two.z) || isnan(one.z) )
    {
        isexit=true;
    }///implement far clipping plane

    if(three.z > 1 && two.z > 1 && one.z > 1)
    {
        isexit=true;
    }///yay, a far clipping plane

    if(isnan (one.y) || isnan(two.y) || isnan(three.y) || isnan(one.x) || isnan(two.x) || isnan(three.x) || isnan(one.z)|| isnan(two.z)|| isnan(three.z))
    {

        isexit=true;

    }

    if((out_of_bounds(one.x, 0, array_width[0]) && out_of_bounds(two.x, 0, array_width[0]) && out_of_bounds(three.x, 0, array_width[0])) || (out_of_bounds(one.y, 0, array_height[0]) && out_of_bounds(two.y, 0, array_height[0]) && out_of_bounds(three.y, 0, array_height[0])))
    {

        isexit=true;

    }


    int y1;
    y1 = round(one.y);
    int y2;
    y2 = round(two.y);
    int y3;
    y3 = round(three.y);

    int x1;
    x1 = round(one.x);
    int x2;
    x2 = round(two.x);
    int x3;
    x3 = round(three.x);

    int miny;
    miny=min3(y1, y2, y3);
    int maxy;
    maxy=max3(y1, y2, y3);
    int minx;
    minx=min3(x1, x2, x3);
    int maxx;
    maxx=max3(x1, x2, x3);

    //__global ulong *tpf=&screena[miny*array_width[0] + minx];
    //prefetch(tpf, abs(maxx-minx)*abs(maxy-miny));

    ///PASS IN ROTATED CAMERA COORDINATES

    float4 camera=c_pos[0];

    //float4 ncamera;
    //ncamera=normalize(camera);

    //float angle = dot(nlight, ncamera);

    float area;
    area=form(x1, y1, x2, y2, x3, y3, 0, 0, 0);



    if(abs(maxx-minx) > array_width[0] || abs(maxy-miny) > array_height[0])
    {
        isexit=true;
    }


    int ndrawn=0;




    minx=max(0, minx);
    maxx=min((int)array_width[0]-1, maxx);

    miny=max(0, miny);
    maxy=min((int)array_height[0]-1, maxy);

    if(isexit){
        return;
    }
    if(!isexit)
    {

        ///project depth on to depth buffer

        for(int x=minx; x<maxx; x++)
        {
            for(int y=miny; y<maxy; y++)
            {


                //if(x < 0 || x >= array_width[0] || y < 0 || y >= array_height[0] ){
                //    continue;
                //}



                float fx, fy;

                fx=(float)x;
                fy=(float)y;



                float s1=fabs(((x2*y-x*y2)+(x3*y2-x2*y3)+(x*y3-x3*y))/2.0) + fabs(((x*y1-x1*y)+(x3*y-x*y3)+(x1*y3-x3*y1))/2.0) + fabs(((x2*y1-x1*y2)+(x*y2-x2*y)+(x1*y-x*y1))/2.0); /// Work out the area of the triangle. If anyone has a better method than the following, i will be happy to implement it

                if(s1 < area + 0.001 && s1 > area - 0.001) /// DO NOT USE 0.0001 THIS CAUSES HOELELELELEL

                {


                    int2 coord;
                    coord.x=x;
                    coord.y=y;



                    float f1=one.z;
                    float f2=two.z;
                    float f3=three.z;


                    float rconstant=(x2*y3+x1*(y2-y3)-x3*y2+(x3-x2)*y1);
                    if(rconstant==0)
                    {
                        continue;
                    }

                    float A=native_divide((f2*y3+f1*(y2-y3)-f3*y2+(f3-f2)*y1),rconstant);
                    float B=native_divide(-(f2*x3+f1*(x2-x3)-f3*x2+(f3-f2)*x1),rconstant);
                    float C=f1-A*x1 - B*y1;

                    float cdepth=(A*x + B*y + C);

                    if(cdepth<depth_cutoff)
                    {
                        continue;
                    }

                    /*if(isnan(cdepth))
                    {

                        continue;

                    }*/


                    uint mydepth=floor(cdepth*(mulint));

                    __global uint *ft=&screena[y*array_width[0] + x];

                    if(floor(cdepth*mulint)>*ft){
                        continue;
                    }


                    uint p=atom_min(ft, mydepth); ///atomic which one is closest to screen


                    if(floor(cdepth*(mulint)) < p)
                    {
                        ndrawn++;
                    }


                }


            }
        }
    }

    //barrier(CLK_GLOBAL_MEM_FENCE);


    ///split kernel here
    /*if(isexit)
    {
        ids[i]=4;

    }
    if(ndrawn==0)
    {
        ids[i]=4;

    }*/
    if(!isexit && ndrawn>0){

        uint id=atom_inc(dcount);
        idbuffer[id]=i;

    }
    }
    //ids[i]=ndrawn;// i + 1
    //int p=0;


    //atomic_inc(dcount);v

}

float4 rotate_3(float4 pos, float4 angle, float4 about){

        float4 ret;
        ret.x=native_cos(angle.y)*(native_sin(angle.z)+native_cos(angle.z)*(pos.x-about.x)) - native_sin(angle.y)*(pos.z-about.z);
        ret.y=native_sin(angle.x)*(native_cos(angle.y)*(pos.z-about.z)+native_sin(angle.y)*(native_sin(angle.z)*(pos.y-about.y)+native_cos(angle.z)*(pos.x-about.x)))+native_cos(angle.x)*(native_cos(angle.z)*(pos.y-about.y)-native_sin(angle.z)*(pos.x-about.x));
        ret.z=native_cos(angle.x)*(native_cos(angle.y)*(pos.z-about.z)+native_sin(angle.y)*(native_sin(angle.z)*(pos.y-about.y)+native_cos(angle.z)*(pos.x-about.x)))-native_sin(angle.x)*(native_cos(angle.z)*(pos.y-about.y)-native_sin(angle.z)*(pos.x-about.x));
        ret.w=0;
        return ret;

}

__kernel void part3(__global float4 *c_pos, __global float4 *c_rot,  __global unsigned int* array_width, __global unsigned int* array_height, __global float4 *stuff, __global uint *screena, __global float4 *lights, __global uint *lightnum, __global uint *ids, __global uint *tcount, __global __write_only image2d_t light_normals, __global __write_only image2d_t uvw_coords, __global __write_only image2d_t depth_vanilla, __global uint *dcount, __global uint *idbuffer)
{
    ///kernels need to do more. first kernel maybe
    //if(!isexit)
    //{

    //while(1){}

    uint i=get_global_id(0);
    if(i>=*dcount){
        return;
    }
    i=idbuffer[i];

    //i=floor(i/4.0f);
    //bool tesselate=false;
    //if(ids[i]<10){
        //tesselate=true;
        //return;
    //}

    if(i<0)
    {
       return;
    }
    if(i >= *tcount)
    {
        return;
    }

    //uint lid=get_global_id(0) % 4;


    ///tesselate everything?

    //if(ids[i]!=0)
    {

        uint location=i*12;

        float4 position[3];
        position[0]=stuff[location + 0];
        position[1]=stuff[location + 1];
        position[2]=stuff[location + 2];

        float4 normal[3];
        normal[0]=stuff[location + 3];
        normal[1]=stuff[location + 4];
        normal[2]=stuff[location + 5];

        float4 colour[3];
        colour[0]=stuff[location + 6];
        colour[1]=stuff[location + 7];
        colour[2]=stuff[location + 8];

        float4 uvw[3];
        uvw[0]=stuff[location + 9];
        uvw[1]=stuff[location + 10];
        uvw[2]=stuff[location + 11];


        ///if(gidbuffer!=dontdraw)
        //ids[i]=0;





        float4 one, two, three; ///should have used an array


        float prez1, prez2, prez3;



        //float depth_cutoff=0.1;


        one.x=native_cos(c_rot[0].y)*(native_sin(c_rot[0].z)+native_cos(c_rot[0].z)*(position[0].x-c_pos[0].x)) - native_sin(c_rot[0].y)*(position[0].z-c_pos[0].z);
        one.y=native_sin(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[0].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[0].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[0].x-c_pos[0].x)))+native_cos(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[0].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[0].x-c_pos[0].x));
        prez1=one.z=native_cos(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[0].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[0].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[0].x-c_pos[0].x)))-native_sin(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[0].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[0].x-c_pos[0].x));

        float rx;
        rx=one.x * (700.0f/one.z);
        float ry;
        ry=one.y * (700.0f/one.z);

        rx+=array_width[0]/2.0f;
        ry+=array_height[0]/2.0f;


        one.x=rx;
        one.y=ry;
        one.z=dcalc(one.z);






        two.x=native_cos(c_rot[0].y)*(native_sin(c_rot[0].z)+native_cos(c_rot[0].z)*(position[1].x-c_pos[0].x)) - native_sin(c_rot[0].y)*(position[1].z-c_pos[0].z);
        two.y=native_sin(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[1].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[1].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[1].x-c_pos[0].x)))+native_cos(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[1].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[1].x-c_pos[0].x));
        prez2=two.z=native_cos(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[1].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[1].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[1].x-c_pos[0].x)))-native_sin(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[1].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[1].x-c_pos[0].x));


        rx=two.x * (700.0f/two.z);
        ry=two.y * (700.0f/two.z);

        rx+=array_width[0]/2.0f;
        ry+=array_height[0]/2.0f;


        two.x=rx;
        two.y=ry;
        two.z=dcalc(two.z);





        three.x=native_cos(c_rot[0].y)*(native_sin(c_rot[0].z)+native_cos(c_rot[0].z)*(position[2].x-c_pos[0].x)) - native_sin(c_rot[0].y)*(position[2].z-c_pos[0].z);
        three.y=native_sin(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[2].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[2].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[2].x-c_pos[0].x)))+native_cos(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[2].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[2].x-c_pos[0].x));
        prez3=three.z=native_cos(c_rot[0].x)*(native_cos(c_rot[0].y)*(position[2].z-c_pos[0].z)+native_sin(c_rot[0].y)*(native_sin(c_rot[0].z)*(position[2].y-c_pos[0].y)+native_cos(c_rot[0].z)*(position[2].x-c_pos[0].x)))-native_sin(c_rot[0].x)*(native_cos(c_rot[0].z)*(position[2].y-c_pos[0].y)-native_sin(c_rot[0].z)*(position[2].x-c_pos[0].x));


        rx=three.x * (700.0f/three.z);
        ry=three.y * (700.0f/three.z);

        rx+=array_width[0]/2.0f;
        ry+=array_height[0]/2.0f;


        three.x=rx;
        three.y=ry;
        three.z=dcalc(three.z);


        int y1;
        y1 = round(one.y);
        int y2;
        y2 = round(two.y);
        int y3;
        y3 = round(three.y);

        int x1;
        x1 = round(one.x);
        int x2;
        x2 = round(two.x);
        int x3;
        x3 = round(three.x);

        int miny;
        miny=min3(y1, y2, y3);
        int maxy;
        maxy=max3(y1, y2, y3);
        int minx;
        minx=min3(x1, x2, x3);
        int maxx;
        maxx=max3(x1, x2, x3);


        minx=max(0, minx);
        maxx=min((int)array_width[0]-1, maxx);

        miny=max(0, miny);
        maxy=min((int)array_height[0]-1, maxy);




        uint ndrawn=0;
        float area;
        area=form(x1, y1, x2, y2, x3, y3, 0, 0, 0);
        for(int x=minx; x<maxx; x = x + 1)
        {
            for(int y=miny; y<maxy; y = y + 1)
            {


                //if(x < 0 || x >= array_width[0] || y < 0 || y >= array_height[0] ){
                //    continue;
                //}


                float fx, fy;

                fx=(float)x;
                fy=(float)y;



                float s1=fabs(((x2*y-x*y2)+(x3*y2-x2*y3)+(x*y3-x3*y))/2.0) + fabs(((x*y1-x1*y)+(x3*y-x*y3)+(x1*y3-x3*y1))/2.0) + fabs(((x2*y1-x1*y2)+(x*y2-x2*y)+(x1*y-x*y1))/2.0);


                if(s1 < area + 0.001 && s1 > area - 0.001) /// DO NOT USE 0.0001 THIS CAUSES HOELELELELEL

                {



                    int2 coord;
                    coord.x=x;
                    coord.y=y;



                    float f1=one.z;
                    float f2=two.z;
                    float f3=three.z;


                    float rconstant=(x2*y3+x1*(y2-y3)-x3*y2+(x3-x2)*y1);
                    if(rconstant==0)
                    {
                        continue;
                    }




                    float A=native_divide((f2*y3+f1*(y2-y3)-f3*y2+(f3-f2)*y1),rconstant);
                    float B=native_divide(-(f2*x3+f1*(x2-x3)-f3*x2+(f3-f2)*x1),rconstant);
                    float C=f1-A*x1 - B*y1;

                    float cdepth=(A*x + B*y + C);

                    if(isnan(cdepth))
                    {
                        continue;
                    }


                    uint cordepth=screena[y*array_width[0] + x];


                    ///If you are not the depth on the depth buffer, go away
                    if(cordepth>round(cdepth*(mulint)) + 50 || cordepth < round(cdepth*(mulint)) - 50)
                    {
                        //float4 temporary={0,0,0,0};
                        //write_imagef(depth_vanilla, coord, temporary);
                        continue;
                    }





                    f1=1.0/prez1;
                    f2=1.0/prez2;
                    f3=1.0/prez3;

                    A=native_divide((f2*y3+f1*(y2-y3)-f3*y2+(f3-f2)*y1),rconstant);
                    B=native_divide(-(f2*x3+f1*(x2-x3)-f3*x2+(f3-f2)*x1),rconstant);

                    C=f1-A*x1 - B*y1;

                    float ldepth=(A*x + B*y + C);


                    /*f1=prez1;
                    f2=prez2;
                    f3=prez3;

                    A=native_divide((f2*y3+f1*(y2-y3)-f3*y2+(f3-f2)*y1),rconstant);
                    B=native_divide(-(f2*x3+f1*(x2-x3)-f3*x2+(f3-f2)*x1),rconstant);

                    C=f1-A*x1 - B*y1;

                    float ldepth1=(A*x + B*y + C);*/


                    f1=uvw[0].x/prez1;
                    f2=uvw[1].x/prez2;
                    f3=uvw[2].x/prez3;

                    A=native_divide((f2*y3+f1*(y2-y3)-f3*y2+(f3-f2)*y1),rconstant);
                    B=native_divide(-(f2*x3+f1*(x2-x3)-f3*x2+(f3-f2)*x1),rconstant);

                    C=f1-A*x1 - B*y1;

                    float uvwx=(A*x + B*y + C);



                    f1=uvw[0].y/prez1;
                    f2=uvw[1].y/prez2;
                    f3=uvw[2].y/prez3;

                    A=native_divide((f2*y3+f1*(y2-y3)-f3*y2+(f3-f2)*y1),rconstant);
                    B=native_divide(-(f2*x3+f1*(x2-x3)-f3*x2+(f3-f2)*x1),rconstant);

                    C=f1-A*x1 - B*y1;

                    float uvwy=(A*x + B*y + C);




                    float4 coord4;
                    coord4.x=(uvwx*1024*1.0/ldepth); ///need to pass texture width, size
                    coord4.y=(uvwy*1024*1.0/ldepth);
                    coord4.z=uvw[0].z; ///all vertex z coordinates are texture ID, should be the same
                    coord4.w=0;

                    int2 coords= {x, y};

                    write_imagef(uvw_coords, coords, coord4);





                    f1=normal[0].z;
                    f2=normal[1].z;
                    f3=normal[2].z;


                    A=native_divide((f2*y3+f1*(y2-y3)-f3*y2+(f3-f2)*y1),rconstant);
                    B=native_divide(-(f2*x3+f1*(x2-x3)-f3*x2+(f3-f2)*x1),rconstant);

                    C=f1-A*x1 - B*y1;

                    float nz=(A*x + B*y + C);




                    f1=normal[0].x;
                    f2=normal[1].x;
                    f3=normal[2].x;

                    A=native_divide((f2*y3+f1*(y2-y3)-f3*y2+(f3-f2)*y1),rconstant);
                    B=native_divide(-(f2*x3+f1*(x2-x3)-f3*x2+(f3-f2)*x1),rconstant);

                    C=f1-A*x1 - B*y1;

                    float nx=(A*x + B*y + C);



                    f1=normal[0].y;
                    f2=normal[1].y;
                    f3=normal[2].y;

                    A=native_divide((f2*y3+f1*(y2-y3)-f3*y2+(f3-f2)*y1),rconstant);
                    B=native_divide(-(f2*x3+f1*(x2-x3)-f3*x2+(f3-f2)*x1),rconstant);

                    C=f1-A*x1 - B*y1;

                    float ny=(A*x + B*y + C);

                    float4 normals= {nx, ny, nz, 0};

                    float4 nought={0,0,0,0};

                    //float4 actual_normal=rotate_3(normals, c_rot[0], nought); ///move to lighting phase?


                    write_imagef(light_normals, coords, fast_normalize(normals));

                    float4 ldepthnull= {1.0/ldepth, 1.0f, 1.0f, 1.0f};
                    int2 temp= {x, y};

                    write_imagef(depth_vanilla, temp, ldepthnull);



                    ndrawn++;

                }

            }
        }

        /*if(ndrawn==0){
            //no pixels drawn
            ids[i]=0;

        }
        //else if(ndrawn<1){
        //    ids[i]=3;
        //}
        else if(ndrawn>200){
            ids[i]=1;
        }
        else
        {

            //atomic_inc(tcount);

            ids[i]=1;

        }

        }*/


        //push good triangles into a buffer


    }
}



__kernel void light(__global float4 *c_pos,__global float4 *c_rot, __global uint *pscreenwidth, __global uint *pscreenheight, __global float4 *lights, __global uint *lightnum, __global __write_only image2d_t screen, __global __read_only image2d_t light_normals, __global __read_only image2d_t uvw, __global uint *screena, __global __read_only image3d_t tex, __global __read_only image2d_t depth_vanilla)
{
///yay, write to alpha channel
///cant write to alpha channel, screen is write only in other kernel


    const sampler_t RAWRAWRAWRAWR=      CLK_NORMALIZED_COORDS_FALSE |
                                        CLK_ADDRESS_CLAMP_TO_EDGE   |
                                        CLK_FILTER_LINEAR;

    uint screenwidth=*pscreenwidth;
    uint screenheight=*pscreenheight;

    float4 lightaccumulation= {0,0,0,0};

    uint x=get_global_id(0);
    uint y=get_global_id(1);

    if(x >= screenwidth || y >= screenheight)
    {
        return;
    }
    if(x < 0 || y < 0)
    {
        return;
    }
    float4 total= {0,0,0,0};
    float4 ncamera=fast_normalize(*c_pos);

    int2 coord= {x, y};

    screena[x + screenwidth*y]=UINT_MAX;

    float4 pnormals=read_imagef(light_normals, RAWRAWRAWRAWR, coord);
    float4 uvwcoords=read_imagef(uvw, RAWRAWRAWRAWR, coord);
    float4 ldepthnull=read_imagef(depth_vanilla, RAWRAWRAWRAWR, coord);
    float4 nought={0,0,0,0};
    float4 rnormals=rotate_3(pnormals, c_rot[0], nought);

    //write_imagef(depth_vanilla, coord, total);

    float ldepth=ldepthnull.x;

    if(isnan(ldepth))
    {
        return;
    }



    if(ldepthnull.y!=0)
    {

        uint lap=0;
        float4 cap={0,0,0,0};
        for(int m=0; m<*lightnum; m++)
        {

            float4 lpos=lights[m*2];
            float4 lcol=lights[(m*2)+1];
            float4 global_lpos=lights[(m*2)+2];


            //float4 pcoord={ldepth*(x-screenwidth/2.0f)/700.0f, ldepth*(y-screenheight/2.0f)/700.0f, ldepth, 0};

            if(isnan(lpos.z))
            {
                continue;
            }


            float radius=lpos.w;

            lpos.w=0;
            global_lpos.w=0;
            float odepth=lpos.z;

            bool treatd=false;


            float4 corrected_pos= {(ldepth*(x-screenwidth/2.0f)/700.0f), (ldepth*(y-screenheight/2.0f)/700.0f), ldepth, 0};

            float4 c_lightpos= {lpos.z*(lpos.x-screenwidth/2.0f)/700.0f, lpos.z*(lpos.y-screenheight/2.0f)/700.0f, lpos.z, 0};

            float distance2=((c_lightpos.x-corrected_pos.x)*(c_lightpos.x-corrected_pos.x) + (c_lightpos.y-corrected_pos.y)*(c_lightpos.y-corrected_pos.y) + ((corrected_pos.z-c_lightpos.z))*((corrected_pos.z-c_lightpos.z))); ///distance makes it dimmer? ///add depth + 10 to lpos.z, add

            if(distance2>radius*radius)
            {
                continue;
            }

            cap+=lcol;

            lap++;


            float distance=sqrt(distance2);





            float4 nlpos=fast_normalize(global_lpos);

            float rough=dot(fast_normalize(c_lightpos-corrected_pos), fast_normalize(rnormals));

            float4 camera_look={0,0,0,0};

            float4 added=(c_lightpos-corrected_pos)+-corrected_pos;
            float4 halflight=(((c_lightpos-corrected_pos)+-corrected_pos)/((fast_length(added))));





            rough+=rabs(rough);

            //rough=rabs(rough);

           // float4 nought={0,0,0,0};

            float aleun=max(0.0, dot(rnormals, halflight)); ///specular lighting failure is probably due to using rotated normals


            float spec_light=0;

            float specular_intensity=0.5f;
            float diffuse_intensity=0.5f;

            rough*=diffuse_intensity;

            float roughness=0.6;



            //aleun=-1.1;
            //aleun*=100;

            if(aleun > 0.0)
            {


                float expt=-native_divide((1-aleun*aleun),(aleun*aleun*roughness*roughness));


                spec_light=specular_intensity*exp(expt);

                spec_light=native_divide(spec_light, (3.14159*roughness*roughness*aleun*aleun*aleun*aleun));

                //spec_light=pow(dot(pnormals, halflight), 1);
                //spec_light=1;

                float beckmann=spec_light;

                float fresnel=pow(1+dot(ncamera, pnormals), 1);

                //float gat=min3(1.0f, (2.0*aleun*dot(ncamera, pnormals))/(dot(ncamera, halflight)), (2.0*aleun*dot(nlpos, pnormals))/(dot(ncamera, halflight))); ///geometric attenuation term

                //spec_light=(beckmann*fresnel*1.0f)/dot(ncamera, pnormals);


            }

            else
            {
                spec_light=0;
            }



            //float specmod=pow(aleun, 10)*specular_intensity;
            float specmod=pow(spec_light, 1)*specular_intensity;

            //specmod=0;

            if(rough>1){
                rough=1;
            }
            if(specmod>1){
                specmod=1;
            }

            if(rough<0 && specmod>0){
                specmod=0;
            }

            //float roughspec=(rough + specmod);

            float roughspec=rough;

            if(roughspec>1)
            {
                //roughspec=1;
            }
            if(roughspec<0)
            {
                roughspec=0;
            }

            float fraction=native_divide(radius,distance);
            fraction*=fraction;
            //float fraction=1.0f;

            if(fraction>1.5f)
            {
                fraction=1.5f;
            }

            fraction*=(1.0f-max(0.0f, distance/radius)); ///something not quite right.


            //total.x=lcol.x*roughspec*fraction + total.x;
            //total.y=lcol.y*roughspec*fraction + total.y;
            //total.z=lcol.z*roughspec*fraction + total.z;
            //rough=0;
            total=lcol*rough*fraction + total + specmod*255.0f;
            ///add spec highlights and do light * spec_highlight


        }

        if(lap==0){
            //return;
        }
        total+=50;
        total/=255.0f; ///divide by largest colour channel?
        //total/=cap;

        //total/=(float)lap;

        float4 colour=read_imagef(tex, RAWRAWRAWRAWR, uvwcoords);



        //float4 colour={255,60,10,0};

        //float dammit=ldepthnull.y;

        //float4 colmod=((total*255.0f+colour)/2.0f)*fast_length(total)/255.0f;
        float4 colmod=total*colour/255.0f;

        if(colmod.x>1)
        {
            colmod.x=1;
        }
        if(colmod.y>1)
        {
            colmod.y=1;
        }
        if(colmod.z>1)
        {
            colmod.z=1;
        }
        if(colmod.x<0)
        {
            colmod.x=0;
        }
        if(colmod.y<0)
        {
            colmod.y=0;
        }
        if(colmod.z<0)
        {
            colmod.z=0;
        }

        int2 coord3= {x,y};



        ///FLIP COORDINATES DUE TO NOW USING OPENGL
        coord3.y=screenheight-coord3.y;
        write_imagef(screen, coord3, colmod);

    }


}
